Skip to content

NanoVDB: support single-space device buffers in GridHandle (CUDA) - #2288

Open
harrism wants to merge 4 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-gridhandle-single-space
Open

NanoVDB: support single-space device buffers in GridHandle (CUDA)#2288
harrism wants to merge 4 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-gridhandle-single-space

Conversation

@harrism

@harrism harrism commented Aug 19, 2026

Copy link
Copy Markdown
Contributor

What this PR is, and why

Grids on the GPU currently require a dual-space buffer: a host allocation mirrored by a device allocation, even in pipelines where the host copy is never read — doubling memory for device-resident workflows and pinning allocation to NanoVDB's built-in allocator instead of the application's pool. This PR (step 3 of #2232, part 1 of 3) adds a device-only alternative: GridHandle<cuda::Buffer<std::byte, R>>, a grid owned by a single device allocation made through any injected memory resource (an RMM-style pool, or downstream wrappers like fvdb's Torch caching allocator — the consumers the #2232 seam work exists for). Host mirror: gone. Host-facing APIs on such a handle: compile-time errors by design.

Dual-space handles are untouched by this PR — today they're still the only way to move grids between host and device. That changes in part 2, which will add explicit cross-space transfers (per-direction copy() between host-space and device-space handles, currently a named compile error). Once those cover the dual buffers' use cases, the dual buffers themselves will be deprecated and then removed: the end state is one buffer template class, with on type per memory space and explicit transfers between them.

What

  • A new (detected) hasDeviceSingle buffer trait selects a constructor in cuda/GridHandle.cuh that parses grid metadata through the device — validated header walk, metadata scratch through the buffer's own resource, kernel, and readback all ordered on the buffer's retained stream.
  • Device accessors work; host paths are compile-time errors. deviceData()/deviceGrid<T>() are available; data(), grid<T>(), gridData(), gridMetaData() are SFINAE-removed; the read/write I/O members and splitGrids/mergeGrids stay addressable (the python bindings take write's address via overload_cast) and instead fail with explanatory static_asserts when instantiated for a single-space handle. The python module (PyGridHandle.h bindings included) is compiled as part of the local gate.
  • copy() dispatches on the traits: host buffers keep the memcpy path; a single-space handle deep-copies device-to-device through its own resource on the retained stream and reuses the host-resident metadata instead of re-parsing. cuda::Buffer gains ElementType/ResourceType, resource(), and a no-arg copy() for stream-ordered resources (= copy(stream())).
  • Hardening shared with the legacy path: the device parse now validates the entire grid chain (per-header bounds checks) before launching the metadata kernel, so truncated buffers and forged mGridCount/mGridSize headers throw instead of reading out of bounds — the pre-existing dual-space parse had the same hole and now uses the same helper. Its scratch runs on MallocResource, so the long-standing GridHandle<DeviceBuffer> device parse keeps working on devices without memory-pool support. Six raw allocation sites in cuda/GridHandle.cuh (one an unchecked cudaMalloc) are replaced with resource-aware buffers.
  • Host accessibility is a resource property: the new is_host_accessible_resource trait detects a HOST_ACCESSIBLE marker (on PinnedResource, forwarded through ResourceRef/AsyncFromSync). A pinned-resource cuda::Buffer handle is a named compile error for now — GridHandle's host paths need an allocation interface cuda::Buffer doesn't yet provide (the create mapping, scheduled for the step-3 completion PR) — rather than a half-working host path.
  • Compatibility: both new trait members are detected, not required — BufferHasDeviceSingle/BufferHasHostSingle default to false when a BufferTraits specialization omits them, so pre-existing specializations in tree (PoolBuffer examples, DeviceBuffer, UnifiedBuffer) and out of tree compile unchanged. Every existing call site of the re-gated members was traced; none changes overload resolution.

Tests

Seven new tests + compile-time classification asserts in TestBuffer.cu: device meta parse and typed/wrong-type deviceGrid; exact allocation accounting through a counting resource (grid bytes + meta scratch, freed via reset()); deep D2D copy with byte comparison; empty-handle and empty-copy; invalid-grid throw with leak check; synchronous-resource (MallocResource) construction and copy; multi-grid parse with a stream-recording resource proving every allocation lands on the retained (non-blocking) stream; forged-mGridCount rejection with an unpoisoned CUDA context.

Verification

  • Zero warnings under the default --Werror=all-warnings; full build (tests, tools, examples) at CMAKE_CUDA_ARCHITECTURES=80; 7/7 ctests on a GPU runner, plus g++ -fsyntax-only on GridHandle.h proving the header stays CUDA-free.
  • Reviewed by an agent audit (consumer trace over every call site, two-phase-lookup check, stream/exception-path analysis) and four Codex review rounds — findings included the pinned-resource misclassification, the chain-validation P1, and a trait-query regression, all addressed; final round clean with the reviewer independently building and running the tests.

🤖 Generated with Claude Code

@harrism
harrism requested a review from kmuseth as a code owner August 19, 2026 23:24
@harrism
harrism force-pushed the nanovdb-gridhandle-single-space branch 4 times, most recently from e270bf0 to 2380d21 Compare August 20, 2026 01:15
@swahtz swahtz added the nanovdb label Aug 20, 2026
@harrism
harrism force-pushed the nanovdb-gridhandle-single-space branch from 2380d21 to e705df9 Compare August 20, 2026 02:01
GridHandle can now own a nanovdb::cuda::Buffer<std::byte, R>: a new
hasDeviceSingle buffer trait selects a constructor (implemented in
cuda/GridHandle.cuh) that parses the grid metadata through the device,
with every operation -- the validated header walk, metadata scratch
allocation through the buffer's resource, the copy kernel and the
readback -- ordered on the buffer's retained stream.
deviceData()/deviceGrid() work on such handles. The host accessors
(data, grid, gridData, gridMetaData) are SFINAE-removed for them,
since the handle owns no host-readable bytes; the read/write I/O
members and splitGrids/mergeGrids stay addressable (the python
bindings take write's address via overload_cast) and fail instead
with an explanatory static_assert when instantiated for a
single-space handle.

The device parse validates the whole grid chain (per-header bounds
checks against the allocation) before launching the metadata kernel,
so truncated buffers and forged mGridCount/mGridSize headers are
rejected with an exception instead of an out-of-bounds device read;
the pre-existing dual-space parse had the same hole and now shares the
validation helper. Its metadata scratch uses MallocResource so that
long-standing path keeps working on devices without memory-pool
support, and the dirty-flag scratch in splitGridHandles and
mergeGridHandles is a resource-aware buffer as well -- together
replacing the file's six raw cudaMalloc/mallocAsync sites, one of
which was missing its cudaCheck.

copy() dispatches on the traits: host buffers keep the memcpy path,
single-space buffers deep-copy device-to-device through their own
resource on the retained stream and reuse the host-resident metadata
instead of re-parsing. Buffer gains ElementType/ResourceType typedefs,
a resource() accessor, and a no-argument copy() for stream-ordered
resources that orders the copy on the retained stream.

Host accessibility is a property of the resource, not the element
type: the new is_host_accessible_resource trait detects a
HOST_ACCESSIBLE marker (declared by PinnedResource and forwarded by
ResourceRef and AsyncFromSync). A cuda::Buffer over a host-accessible
resource is rejected at handle scope with a named error for now --
GridHandle's host paths require an allocation interface cuda::Buffer
does not yet provide -- and non-byte element types fail loudly in the
single-space constructor. Both hasDeviceSingle and hasHostSingle are
detected rather than required, so pre-existing BufferTraits
specializations in and out of tree compile unchanged.

Part of AcademySoftwareFoundation#2232 (step 3).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
@harrism
harrism force-pushed the nanovdb-gridhandle-single-space branch from e705df9 to c927ae7 Compare August 20, 2026 05:06

@swahtz swahtz left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reviewed with a VBM follow-up as part of #2232 in mind. The design lands the right conventions — the detected traits, the named compile errors instead of half-working host paths, and never requiring default-constructible buffers are all things the next handle types can build on directly. Two issues inline: a correctness regression where the hardened chain validation now throws on addBlindData output from multi-grid sources, and the O(N) stream syncs the validation adds to the long-standing GridHandle constructor (with two possible shapes for fixing it). The rest are non-blocking polish.

Comment thread nanovdb/nanovdb/cuda/GridHandle.cuh Outdated
/// the device-side metadata walk that follows can never read out of
/// bounds from a truncated buffer or a forged header.
/// @return the validated grid count
inline uint32_t validGridChainCount(const GridData *d_head, uint64_t bytes, cudaStream_t stream)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🚩 issue: I traced the consumers of the device-parse path, and I believe this new validation breaks tools::cuda::addBlindData for grids that came out of a multi-grid handle. addBlindData copies the source grid's GridData header verbatim (including mGridIndex/mGridCount) into its new single-grid buffer and never normalizes them before constructing GridHandle<BufferT> (AddBlindData.cuh:131, a device-only buffer, so it takes this path). With a source grid from handle.deviceGrid<T>(n) of a multi-grid handle: for n > 0 the first iteration throws "inconsistent grid index/count" (mGridIndex != 0), and for n == 0 with mGridCount > 1 the walk runs past the single-grid buffer and throws "grid chain exceeds the device buffer". The old parse accepted these buffers (albeit with OOB metadata reads — which is exactly the hole this validation fixes).

I think the fix belongs in addBlindData rather than in weakening the validation: normalize the header to index 0 / count 1 before constructing the handle, the way splitGridHandles does with detail::updateGridCount<<<1,1>>>. Might also be worth a quick audit for any other tool that constructs a device-parse handle from a copied header.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in 5ea0b9e — normalized to index 0 / count 1 inside the existing header-fixup kernel (ahead of the tool's checksum recompute, which covers the change). Your audit hunch was right: indexToGrid copies the source header the same way (*dstGrid.data() = *srcGrid.data() and then patches only mGridType/mData1), fixed in the same commit. The other handle-returning tools are clean — the topology ops and MeshToGrid normalize in TopologyBuilder, and the point builders initialize fresh headers via GridData::init. Regression tests run grids 0 and 1 of a merged two-grid handle through both tools; verified failing before the fix with exactly the two throws you predicted.

Comment thread nanovdb/nanovdb/cuda/GridHandle.cuh Outdated
mMetaData.resize(tmp.mGridCount);
cudaCheck(cudaMemcpy(mMetaData.data(), d_metaData,tmp.mGridCount*sizeof(GridHandleMetaData), cudaMemcpyDeviceToHost));
cudaCheck(cudaFree(d_metaData));
const uint32_t count = cuda::detail::validGridChainCount(d_data, mBuffer.size(), cudaStream_t(0));

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🚩 issue: This long-standing constructor previously did a single header read; it now does one cudaMemcpyAsync + full cudaStreamSynchronize per grid header inside validGridChainCount. cuda::mergeGridHandles ends by constructing exactly this handle, so merging hundreds of grids now issues hundreds of serialized 672-byte D2H copies, each draining the stream, where the old code did one memcpy plus one kernel.

Two shapes that would fix it, either of which I'd be happy with: (a) since the validation walk already copies every full GridData header to the host, accumulate mMetaData directly in the walk — that deletes the scratch allocation, the cpyGridHandleMeta<<<1,1>>> launch, the second D2H copy, and the final sync in both constructors, so the validation replaces the old parse machinery instead of adding to it (the round trips stay O(N), but the whole second phase disappears); or (b) move the validation into a single device-side kernel with one status readback, restoring O(1) round trips for the multi-grid case.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in 0f4395d with shape (b) plus the deletion from (a): one host read of the head header (still needed to size the metadata scratch and bound a forged mGridCount before allocating), then a single kernel that validates the chain and fills the metadata scratch, one readback, one sync. So the whole second phase — the cpyGridHandleMeta launch, its D2H copy, and the final sync — is gone from both constructors, and constructing the handle mergeGridHandles returns costs two round trips regardless of grid count.

Comment thread nanovdb/nanovdb/GridHandle.h Outdated
/// or if the template parameter does not match the specified grid.
template<typename ValueT, typename U = BufferT>
typename util::enable_if<BufferHasDeviceSingle<U>::value, const NanoGrid<ValueT>*>::type
deviceGrid(uint32_t n=0) const {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💅 polish: This body is a token-for-token copy of grid() and the dual-space deviceGrid() — the same null/index/gridType checks and util::PtrAdd, differing only in the base pointer (mBuffer.data() vs mBuffer.deviceData()). A private ungated helper taking the base pointer (e.g. gridAt<ValueT>(const void* base, uint32_t n)) would serve all three, so a future change to the lookup contract doesn't need three edits — this inline copy sits far from the two out-of-line definitions and would be the one that gets missed.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in e86e55e — a private gridAt<ValueT>(base, n) now backs grid() and both deviceGrid() families.

@@ -371,14 +462,40 @@ template<typename BufferT>
template <typename OtherBufferT>
inline GridHandle<OtherBufferT> GridHandle<BufferT>::copy(const OtherBufferT& other) const

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 suggestion: Both copy() overloads carry an identical if constexpr guard and multi-line static_assert, and this overload's single-space branch just discards other and forwards to copy<OtherBufferT>() — which re-fires the same assert, so today every misuse emits the diagnostic twice. I believe the guard can live solely in the no-arg copy() with identical diagnostics for every (BufferT, OtherBufferT) combination. This matters for part 2: when cross-space transfers land, the "not supported yet" condition and message get relaxed, and with two copies of the guard the two spellings can silently drift to accepting different type combinations.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in e86e55e — the guard lives solely in the no-arg copy(); the pool overload forwards its single-space branch there uncommented, so a misuse diagnoses once and the two spellings cannot drift when cross-space transfers relax the condition.

Comment thread nanovdb/nanovdb/GridHandle.h Outdated
/// BufferTraits specializations (in or out of tree) that only define
/// hasDeviceDual keep compiling unchanged.
template<typename BufferT, typename = void>
struct BufferHasDeviceSingle { static constexpr bool value = false; };

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 suggestion: Could we move BufferHasDeviceSingle/BufferHasHostSingle next to the BufferTraits primary in HostBuffer.h (or a small shared traits header)? They aren't GridHandle-specific — they're companions to the trait protocol itself — and there's a concrete second consumer lined up: VoxelBlockManagerHandle needs exactly the same dispatch to map deviceFirstLeafID()/deviceJumpMap() onto data() for single-space buffers (the VBM follow-up to this PR under #2232), and it shouldn't have to include GridHandle.h, or duplicate the detectors, to get it.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in e86e55e — both detectors now sit beside the BufferTraits primary in HostBuffer.h, so VoxelBlockManagerHandle gets them without GridHandle.h.

/// (stream-ordered resources), and the (stream, resource, count, noInit)
/// constructor shape.
template<typename T, typename R>
struct BufferTraits<cuda::Buffer<T, R>>

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 suggestion: This @note is effectively the definition of a "single-space buffer concept" — it lists exactly the members a buffer must provide for hasDeviceSingle to be honored. Since other handle types will consume the same contract (VoxelBlockManagerHandle is next, per the #2232 plan), could we phrase it as such — name the concept, and state that any consumer of hasDeviceSingle may rely on exactly this interface? That lets the follow-up PRs cite it instead of reverse-engineering the GridHandle constructor.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in e86e55e — the @note now defines the single-space device-buffer concept and enumerates the exact interface a consumer of hasDeviceSingle may rely on. It also got tighter while writing it down: the metadata scratch now allocates through resource() as a cuda::Buffer, so the constructor shape is no longer part of the contract — a conforming buffer never needs to be constructible by a consumer.

Comment thread nanovdb/nanovdb/cuda/Buffer.h
harrism and others added 3 commits August 20, 2026 07:58
Replace the per-header device-to-host walk in the GridHandle constructors
with a single host read of the head header (which sizes the metadata
scratch) followed by one kernel that validates every header and fills the
metadata scratch in the same pass, plus one readback and one
synchronization regardless of the grid count. The previous walk issued one
synchronizing copy per header, so constructing the handle that
cuda::mergeGridHandles returns cost O(N) stream drains for N grids; it now
costs the same two round trips as a single-grid buffer, and the separate
metadata kernel launch, its scratch copy and final synchronization are
gone. Bounds are still checked before every header read, on the host for
the head and on the device for the rest, so a truncated buffer or forged
header can never read out of bounds.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…utput

Both tools copy the source grid's header verbatim into a new single-grid
buffer, so a source grid taken from a multi-grid handle left a stale
mGridIndex/mGridCount in the output. The returned handle's metadata parse
then walked a grid chain that is not there: past the end of the buffer for
index 0 of a multi-grid source, or starting at a nonzero index. The output
header is now normalized to index 0 / count 1 in the same kernels that
already patch its other fields, before each tool's existing checksum
update. The other handle-returning GPU tools were audited: the topology
ops and MeshToGrid normalize in TopologyBuilder, and the point builders
initialize fresh headers via GridData::init.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…lookups

BufferHasDeviceSingle/BufferHasHostSingle move next to the BufferTraits
primary in HostBuffer.h: they are companions to the trait protocol, not
GridHandle-specific, so other handle types can consume them without
including GridHandle.h. The three identical grid-lookup bodies behind
grid() and both deviceGrid() families collapse into one private gridAt()
helper, and the supported-combination guard for GridHandle::copy() now
lives solely in the no-arg overload, so the two spellings cannot drift
apart when cross-space transfers relax it. cuda::Buffer documents the
single-space device-buffer concept on its BufferTraits specialization --
naming exactly the interface a consumer of hasDeviceSingle may rely on --
and gains a rebind alias for generic code that constructs sibling buffers.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants